#include <CL/opencl.h>
#include <fstream>
#include <iostream>
#include <conio.h>
#include <time.h>

const char* source_file = "test.cl";

void *source_a;
void *source_b;
float *dest;
float *result;

unsigned long time_start;
unsigned long time_end;
unsigned long time_gpu;
unsigned long time_cpu;

int num_elements = 10485760;

cl_context dev_context;
cl_device_id device;
cl_command_queue command_que;
cl_program program;
cl_kernel kernel;
cl_platform_id platform;      // OpenCL platform

cl_mem dev_source_a;				// OpenCL device source buffer A
cl_mem dev_source_b;				// OpenCL device source buffer B 
cl_mem dev_dest;					// OpenCL device destination buffer 

size_t global_work_size;
size_t local_work_size;

size_t kernel_length;			// Byte size of kernel code

void Add(const float* data1, const float* data2, float* result, int num );
void Release (int exit_code);

// Round Up Division function
size_t shrRoundUp(int group_size, int global_size) 
{
    int r = global_size % group_size;
    if(r == 0) 
    {
        return global_size;
    } else 
    {
        return global_size + group_size - r;
    }
}


char* oclLoadProgSource(const char* cFilename, const char* cPreamble, size_t* szFinalLength)
{
    // locals 
    FILE* pFileStream = NULL;
    size_t szSourceLength;

    // open the OpenCL source code file
    if(fopen_s(&pFileStream, cFilename, "rb") != 0) 
    {       
        return NULL;
    }


    size_t szPreambleLength = strlen(cPreamble);

    // get the length of the source code
    fseek(pFileStream, 0, SEEK_END); 
    szSourceLength = ftell(pFileStream);
    fseek(pFileStream, 0, SEEK_SET); 

    // allocate a buffer for the source code string and read it in
    char* cSourceString = (char *)malloc(szSourceLength + szPreambleLength + 1); 
    memcpy(cSourceString, cPreamble, szPreambleLength);
    if (fread((cSourceString) + szPreambleLength, szSourceLength, 1, pFileStream) != 1)
    {
        fclose(pFileStream);
        free(cSourceString);
        return 0;
    }

    // close the file and return the total length of the combined (preamble + source) string
    fclose(pFileStream);
    if(szFinalLength != 0)
    {
        *szFinalLength = szSourceLength + szPreambleLength;
    }
    cSourceString[szSourceLength + szPreambleLength] = '\0';

    return cSourceString;
}


bool Compare( const float* reference, const float* data, const unsigned int len ) 
{
    for (int i = 0; i < len; i++)
	{
		if (reference[i] != data[i] )
		{
			return false;
		}
	}

	return true;
}


void FillArray(float* data, int size)
{
    int i; 
    const float scale = 1.0f / (float)RAND_MAX;
    for (i = 0; i < size; ++i) 
    {
        data[i] = scale * rand();
    }
}


int main (int argc, char **argv)
{	
	int error = 0;
	local_work_size = 256;
	global_work_size = shrRoundUp((int)local_work_size, num_elements);  // rounded up to the nearest multiple of the LocalWorkSize

	printf("Global Work Size \t\t= %u\nLocal Work Size \t\t= %u\n# of Work Groups \t\t= %u\n\n", 
           global_work_size, local_work_size, (global_work_size % local_work_size + global_work_size/local_work_size)); 

	source_a	= (void *)malloc(sizeof(cl_float) * global_work_size);
    source_b	= (void *)malloc(sizeof(cl_float) * global_work_size);
    dest		= (float *)malloc(sizeof(cl_float) * global_work_size);
    result		= (float *)malloc(sizeof(cl_float) * num_elements);

	FillArray((float*)source_a, num_elements);
    FillArray((float*)source_b, num_elements);

	error = clGetPlatformIDs(1, &platform, NULL);
	if (error != CL_SUCCESS )
	{
		printf("platform error\n");
		Release(EXIT_FAILURE);
	}

	error = clGetDeviceIDs(platform, CL_DEVICE_TYPE_GPU, 1, &device, NULL);
	if (error != CL_SUCCESS )
	{
		printf("device error\n");
		Release(EXIT_FAILURE);
	}

	dev_context = clCreateContext(0, 1, &device, NULL, NULL, &error);
	if ( error != CL_SUCCESS )
	{	
		printf("dev_context error");
		Release(EXIT_FAILURE);
	}
	
	command_que = clCreateCommandQueue(dev_context, device, 0, &error);
	if (error != CL_SUCCESS)
	{
		printf("command_que error\n");
		Release(error);
	}

	dev_source_a = clCreateBuffer(dev_context, CL_MEM_READ_ONLY, sizeof(cl_float) * global_work_size, NULL, &error);
	if (error != CL_SUCCESS )
	{
		printf("dev_source_a error\n");
		Release(EXIT_FAILURE);
	}

	dev_source_b = clCreateBuffer(dev_context, CL_MEM_READ_ONLY, sizeof(cl_float) * global_work_size, NULL, &error);
	if (error!= CL_SUCCESS)
	{
		printf("dev_source_b error\n");
		Release(EXIT_FAILURE);
	}

	dev_dest = clCreateBuffer(dev_context, CL_MEM_WRITE_ONLY, sizeof(cl_float) * global_work_size, NULL, &error);
	if (error != CL_SUCCESS)
	{
		printf("dev_dest error\n");
		Release(EXIT_FAILURE);
	}

	char *path = oclLoadProgSource(source_file, "", &kernel_length);
	program = clCreateProgramWithSource(dev_context, 1,(const char **)&path, &kernel_length, &error);
	if (error != CL_SUCCESS)
	{
		printf("program error\n");
		Release(EXIT_FAILURE);
	}

	error = clBuildProgram(program, 0, NULL, NULL, NULL, NULL);
	if (error != CL_SUCCESS)
	{
		printf("clBuildProgram error\n");
		Release(EXIT_FAILURE);
	}

	kernel = clCreateKernel(program, "add", &error);
	if (error != CL_SUCCESS)
	{
		printf("kernel error\n");
		Release(EXIT_FAILURE);
	}

	error = clSetKernelArg(kernel, 0, sizeof(cl_mem), (void*)&dev_source_a);
	error |= clSetKernelArg(kernel, 1, sizeof(cl_mem), (void*)&dev_source_b);
	error |= clSetKernelArg(kernel, 2, sizeof(cl_mem), (void*)&dev_dest);
	error |= clSetKernelArg(kernel, 3, sizeof(cl_int), (void*)&num_elements);
	if (error != CL_SUCCESS)
	{
		printf("clSetKernelArg error\n");
		Release(EXIT_FAILURE);
	}

	error = clEnqueueWriteBuffer(command_que, dev_source_a, CL_FALSE, 0, sizeof(cl_float) * global_work_size, source_a, 0, NULL, NULL);
	error |= clEnqueueWriteBuffer(command_que, dev_source_b, CL_FALSE, 0, sizeof(cl_float) * global_work_size, source_b, 0, NULL, NULL);
	if (error != CL_SUCCESS)
	{
		printf("clEnqueueWriteBuffer error\n");
		Release(EXIT_FAILURE);
	}
	
	do {
		error = clEnqueueNDRangeKernel(command_que, kernel, 1, NULL, &global_work_size, NULL, 0, NULL, NULL);

		if (error != CL_SUCCESS)
		{
			printf("clEnqueueNDRangeKernel error\n");
			Release(EXIT_FAILURE);
		}

		
		time_start = clock ();
		error = clEnqueueReadBuffer(command_que, dev_dest, CL_TRUE, 0, sizeof(cl_float) * global_work_size, dest, 0, NULL, NULL);
		time_end = clock ();
		if (error != CL_SUCCESS)
		{
			printf("clEnqueueNDRangeKernel error\n");
			Release(EXIT_FAILURE);
		}
		time_gpu += time_end - time_start;

		time_start = clock ();
		Add((const float*)source_a, (const float*)source_b, (float*)result, num_elements);
		time_end = clock ();
		time_cpu += time_end - time_start;

		bool check = Compare((const float*)result, (const float*)dest, (unsigned int)num_elements );
		printf("%s\n\n", (check == true) ? "PASSED" : "FAILED");
		printf("Time  CPU=%u GPU=%u\n", time_cpu, time_gpu);
		printf("Repeate?\n");
	}
	while (getch() == 'y');
	

	Release(EXIT_SUCCESS);
	return 0;
}

void Release (int error)
{  
	if (kernel)			clReleaseKernel(kernel);  
    if (program)		clReleaseProgram(program);
    if (command_que)	clReleaseCommandQueue(command_que);
    if (dev_context)	clReleaseContext(dev_context);
    if (dev_source_a)	clReleaseMemObject(dev_source_a);
    if (dev_source_b)	clReleaseMemObject(dev_source_b);
    if (dev_dest)		clReleaseMemObject(dev_dest);

	free(source_a); 
    free(source_b);
    free(dest);

	if (error)
	{
		printf("FAIL!\n");
	}
	else
	{
		printf("OK!\n");
	}
	getchar();

	exit (error);
}


void Add(const float* data1, const float* data2, float* result, int num )
{
	int i;
    for (i = 0; i < num; i++) 
    {
        result[i] = data1[i] + data2[i]; 
    }
}